[AMDGPU] Add a new builtin type for image descriptor rsrc #160258

ranapratap55 · 2025-09-23T09:19:39Z

Adding a new builtin type for AMDGPU's image descriptor rsrc data type

This requires for #140210

llvmbot · 2025-09-23T09:20:12Z

@llvm/pr-subscribers-clang-codegen
@llvm/pr-subscribers-backend-amdgpu
@llvm/pr-subscribers-clang

@llvm/pr-subscribers-llvm-ir

Author: Rana Pratap Reddy (ranapratap55)

Changes

Adding a new builtin type for AMDGPU's image descriptor rsrc data type

This requires for #140210

Full diff: https://github.com/llvm/llvm-project/pull/160258.diff

11 Files Affected:

(modified) clang/include/clang/Basic/AMDGPUTypes.def (+7)
(modified) clang/include/clang/Basic/Builtins.def (+1)
(modified) clang/lib/AST/ASTContext.cpp (+4)
(modified) clang/lib/CodeGen/CGDebugInfo.cpp (+8)
(modified) clang/lib/CodeGen/CodeGenTypes.cpp (+4)
(added) clang/test/CodeGen/amdgpu-image-rsrc-type-debug-info.c (+18)
(added) clang/test/CodeGenCXX/amdgpu-image-rsrc-typeinfo.cpp (+7)
(added) clang/test/SemaCXX/amdgpu-image-rsrc.cpp (+21)
(added) clang/test/SemaOpenCL/amdgpu-image-rsrc.cl (+14)
(added) clang/test/SemaOpenMP/amdgpu-image-rsrc.cpp (+12)
(modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+1)

diff --git a/clang/include/clang/Basic/AMDGPUTypes.def b/clang/include/clang/Basic/AMDGPUTypes.def
index d3dff446f9edf..8c0bd73252c50 100644
--- a/clang/include/clang/Basic/AMDGPUTypes.def
+++ b/clang/include/clang/Basic/AMDGPUTypes.def
@@ -20,10 +20,17 @@
   AMDGPU_TYPE(Name, Id, SingletonId, Width, Align)
 #endif
 
+#ifndef AMDGPU_IMAGE_RSRC_TYPE
+#define AMDGPU_IMAGE_RSRC_TYPE(Name, Id, SingletonId) \
+  AMDGPU_TYPE(Name, Id, SingletonId, 256, 256)
+#endif
+
 AMDGPU_OPAQUE_PTR_TYPE("__amdgpu_buffer_rsrc_t", AMDGPUBufferRsrc, AMDGPUBufferRsrcTy, 128, 128, 8)
+AMDGPU_IMAGE_RSRC_TYPE("__amdgpu_image_rsrc_t", AMDGPUImageDescRsrc, AMDGPUImageDescRsrcTy)
 
 AMDGPU_NAMED_BARRIER_TYPE("__amdgpu_named_workgroup_barrier_t", AMDGPUNamedWorkgroupBarrier, AMDGPUNamedWorkgroupBarrierTy, 128, 32, 0)
 
 #undef AMDGPU_TYPE
 #undef AMDGPU_OPAQUE_PTR_TYPE
 #undef AMDGPU_NAMED_BARRIER_TYPE
+#undef AMDGPU_IMAGE_RSRC_TYPE
diff --git a/clang/include/clang/Basic/Builtins.def b/clang/include/clang/Basic/Builtins.def
index 48437c9397570..a91315680f93f 100644
--- a/clang/include/clang/Basic/Builtins.def
+++ b/clang/include/clang/Basic/Builtins.def
@@ -34,6 +34,7 @@
 //  Q -> target builtin type, followed by a character to distinguish the builtin type
 //    Qa -> AArch64 svcount_t builtin type.
 //    Qb -> AMDGPU __amdgpu_buffer_rsrc_t builtin type.
+//    Qc -> AMDGPU __amdgpu_image_desc_t builtin type.
 //  E -> ext_vector, followed by the number of elements and the base type.
 //  X -> _Complex, followed by the base type.
 //  Y -> ptrdiff_t
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 97c59b2ceec2f..7ba1dfed1c0db 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -12580,6 +12580,10 @@ static QualType DecodeTypeFromStr(const char *&Str, const ASTContext &Context,
       Type = Context.AMDGPUBufferRsrcTy;
       break;
     }
+    case 'c': {
+      Type = Context.AMDGPUImageDescRsrcTy;
+      break;
+    }
     default:
       llvm_unreachable("Unexpected target builtin type");
     }
diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp
index 578d09f7971d6..3f5277ab66f55 100644
--- a/clang/lib/CodeGen/CGDebugInfo.cpp
+++ b/clang/lib/CodeGen/CGDebugInfo.cpp
@@ -1020,6 +1020,14 @@ llvm::DIType *CGDebugInfo::CreateType(const BuiltinType *BT) {
           DBuilder.createBasicType(Name, Width, llvm::dwarf::DW_ATE_unsigned); \
     return SingletonId;                                                        \
   }
+#define AMDGPU_IMAGE_RSRC_TYPE(Name, Id, SingletonId)                          \
+  case BuiltinType::Id: {                                                      \
+    if (!SingletonId)                                                          \
+      SingletonId =                                                            \
+          DBuilder.createForwardDecl(llvm::dwarf::DW_TAG_structure_type, Name, \
+                                     TheCU, TheCU->getFile(), 0);              \
+    return SingletonId;                                                        \
+  }
 #include "clang/Basic/AMDGPUTypes.def"
   case BuiltinType::UChar:
   case BuiltinType::Char_U:
diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp
index 3ffe999d01178..e3e44556ce514 100644
--- a/clang/lib/CodeGen/CodeGenTypes.cpp
+++ b/clang/lib/CodeGen/CodeGenTypes.cpp
@@ -581,6 +581,10 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) {
   case BuiltinType::Id:                                                        \
     return llvm::TargetExtType::get(getLLVMContext(), "amdgcn.named.barrier",  \
                                     {}, {Scope});
+#define AMDGPU_IMAGE_RSRC_TYPE(Name, Id, SingletonId)                          \
+  case BuiltinType::Id:                                                        \
+    return llvm::VectorType::get(llvm::Type::getInt32Ty(getLLVMContext()), 8,  \
+                                 false);
 #include "clang/Basic/AMDGPUTypes.def"
 #define HLSL_INTANGIBLE_TYPE(Name, Id, SingletonId) case BuiltinType::Id:
 #include "clang/Basic/HLSLIntangibleTypes.def"
diff --git a/clang/test/CodeGen/amdgpu-image-rsrc-type-debug-info.c b/clang/test/CodeGen/amdgpu-image-rsrc-type-debug-info.c
new file mode 100644
index 0000000000000..0e42420e26322
--- /dev/null
+++ b/clang/test/CodeGen/amdgpu-image-rsrc-type-debug-info.c
@@ -0,0 +1,18 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn -emit-llvm -o - %s -debug-info-kind=limited | FileCheck %s
+
+// CHECK-LABEL: define dso_local void @test_locals(
+// CHECK-SAME: ) #[[ATTR0:[0-9]+]] !dbg [[DBG6:![0-9]+]] {
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK-NEXT:    [[IMG:%.*]] = alloca <8 x i32>, align 32, addrspace(5)
+// CHECK-NEXT:    [[IMG_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[IMG]] to ptr
+// CHECK-NEXT:      #dbg_declare(ptr addrspace(5) [[IMG]], [[META11:![0-9]+]], !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef), [[META14:![0-9]+]])
+// CHECK-NEXT:    [[TMP0:%.*]] = load <8 x i32>, ptr [[IMG_ASCAST]], align 32, !dbg [[DBG15:![0-9]+]]
+// CHECK-NEXT:    ret void, !dbg [[DBG16:![0-9]+]]
+//
+void test_locals(void) {
+  __amdgpu_image_rsrc_t img;
+  (void)img;
+}
+
diff --git a/clang/test/CodeGenCXX/amdgpu-image-rsrc-typeinfo.cpp b/clang/test/CodeGenCXX/amdgpu-image-rsrc-typeinfo.cpp
new file mode 100644
index 0000000000000..d96cf5f35c5b2
--- /dev/null
+++ b/clang/test/CodeGenCXX/amdgpu-image-rsrc-typeinfo.cpp
@@ -0,0 +1,7 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn %s -emit-llvm -o - | FileCheck %s
+namespace std { class type_info; }
+auto &a = typeid(__amdgpu_image_rsrc_t);
+//// NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
+// CHECK: {{.*}}
diff --git a/clang/test/SemaCXX/amdgpu-image-rsrc.cpp b/clang/test/SemaCXX/amdgpu-image-rsrc.cpp
new file mode 100644
index 0000000000000..1a19a94039b5e
--- /dev/null
+++ b/clang/test/SemaCXX/amdgpu-image-rsrc.cpp
@@ -0,0 +1,21 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -fsyntax-only -verify -std=gnu++11 -triple amdgcn -Wno-unused-value %s
+
+void foo() {
+  int n = 1;
+  __amdgpu_image_rsrc_t v = 0;            // expected-error {{cannot initialize a variable of type '__amdgpu_image_rsrc_t' with an rvalue of type 'int'}}
+  static_cast<__amdgpu_image_rsrc_t>(n); // expected-error {{static_cast from 'int' to '__amdgpu_image_rsrc_t' is not allowed}}
+  reinterpret_cast<__amdgpu_image_rsrc_t>(n); // expected-error {{reinterpret_cast from 'int' to '__amdgpu_image_rsrc_t' is not allowed}}
+  (void)(v + v); // expected-error {{invalid operands}}
+  int x(v);      // expected-error {{cannot initialize a variable of type 'int' with an lvalue of type '__amdgpu_image_rsrc_t'}}
+  __amdgpu_image_rsrc_t k;
+}
+
+static_assert(sizeof(__amdgpu_image_rsrc_t) == 32, "size");
+static_assert(alignof(__amdgpu_image_rsrc_t) == 32, "align");
+
+template<class T> void bar(T);
+void use(__amdgpu_image_rsrc_t r) { bar(r); }
+struct S { __amdgpu_image_rsrc_t r; int a; };
+static_assert(sizeof(S) == 64, "struct layout");
diff --git a/clang/test/SemaOpenCL/amdgpu-image-rsrc.cl b/clang/test/SemaOpenCL/amdgpu-image-rsrc.cl
new file mode 100644
index 0000000000000..341ab667ebd06
--- /dev/null
+++ b/clang/test/SemaOpenCL/amdgpu-image-rsrc.cl
@@ -0,0 +1,14 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -verify -cl-std=CL1.2 -triple amdgcn-amd-amdhsa %s
+// RUN: %clang_cc1 -verify -cl-std=CL2.0 -triple amdgcn-amd-amdhsa %s
+
+void f() {
+    int n = 3;
+    __amdgpu_image_rsrc_t v = 0; // expected-error {{initializing '__private __amdgpu_image_rsrc_t' with an expression of incompatible type 'int'}}
+    int k = v;                   // expected-error {{initializing '__private int' with an expression of incompatible type '__private __amdgpu_image_rsrc_t'}}
+    (void)(v + v);               // expected-error {{invalid operands}}
+    __amdgpu_image_rsrc_t r;
+    int *p = (int*)r;            // expected-error {{operand of type '__amdgpu_image_rsrc_t' where arithmetic or pointer type is required}}
+    (void)p;
+}
diff --git a/clang/test/SemaOpenMP/amdgpu-image-rsrc.cpp b/clang/test/SemaOpenMP/amdgpu-image-rsrc.cpp
new file mode 100644
index 0000000000000..91d566be9b8a3
--- /dev/null
+++ b/clang/test/SemaOpenMP/amdgpu-image-rsrc.cpp
@@ -0,0 +1,12 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -Wno-unused-value %s
+
+void foo() {
+#pragma omp target
+  {
+    int n = 5;
+    __amdgpu_image_rsrc_t v = 0; // expected-error {{cannot initialize a variable of type '__amdgpu_image_rsrc_t' with an rvalue of type 'int'}}
+    (void)(v + v);               // expected-error {{invalid operands to binary expression ('__amdgpu_image_rsrc_t' and '__amdgpu_image_rsrc_t'}}
+  }
+}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index afce1fe6af854..d41ce5b64b7cd 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -967,6 +967,7 @@ class AMDGPUDimProps<bits<3> enc, string name, string asmsuffix,
   bits<8> NumGradients = !size(GradientArgs);
 }
 
+def AMDGPUImageDescRsrcTy : LLVMType<v8i32>;
 def AMDGPUDim1D : AMDGPUDimProps<0x0, "1d", "1D", ["s"], []>;
 def AMDGPUDim2D : AMDGPUDimProps<0x1, "2d", "2D", ["s", "t"], []>;
 def AMDGPUDim3D : AMDGPUDimProps<0x2, "3d", "3D", ["s", "t", "r"], []>;

llvmbot · 2025-09-23T09:20:13Z

@llvm/pr-subscribers-debuginfo

Author: Rana Pratap Reddy (ranapratap55)

Changes

Adding a new builtin type for AMDGPU's image descriptor rsrc data type

This requires for #140210

Full diff: https://github.com/llvm/llvm-project/pull/160258.diff

11 Files Affected:

(modified) clang/include/clang/Basic/AMDGPUTypes.def (+7)
(modified) clang/include/clang/Basic/Builtins.def (+1)
(modified) clang/lib/AST/ASTContext.cpp (+4)
(modified) clang/lib/CodeGen/CGDebugInfo.cpp (+8)
(modified) clang/lib/CodeGen/CodeGenTypes.cpp (+4)
(added) clang/test/CodeGen/amdgpu-image-rsrc-type-debug-info.c (+18)
(added) clang/test/CodeGenCXX/amdgpu-image-rsrc-typeinfo.cpp (+7)
(added) clang/test/SemaCXX/amdgpu-image-rsrc.cpp (+21)
(added) clang/test/SemaOpenCL/amdgpu-image-rsrc.cl (+14)
(added) clang/test/SemaOpenMP/amdgpu-image-rsrc.cpp (+12)
(modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+1)

diff --git a/clang/include/clang/Basic/AMDGPUTypes.def b/clang/include/clang/Basic/AMDGPUTypes.def
index d3dff446f9edf..8c0bd73252c50 100644
--- a/clang/include/clang/Basic/AMDGPUTypes.def
+++ b/clang/include/clang/Basic/AMDGPUTypes.def
@@ -20,10 +20,17 @@
   AMDGPU_TYPE(Name, Id, SingletonId, Width, Align)
 #endif
 
+#ifndef AMDGPU_IMAGE_RSRC_TYPE
+#define AMDGPU_IMAGE_RSRC_TYPE(Name, Id, SingletonId) \
+  AMDGPU_TYPE(Name, Id, SingletonId, 256, 256)
+#endif
+
 AMDGPU_OPAQUE_PTR_TYPE("__amdgpu_buffer_rsrc_t", AMDGPUBufferRsrc, AMDGPUBufferRsrcTy, 128, 128, 8)
+AMDGPU_IMAGE_RSRC_TYPE("__amdgpu_image_rsrc_t", AMDGPUImageDescRsrc, AMDGPUImageDescRsrcTy)
 
 AMDGPU_NAMED_BARRIER_TYPE("__amdgpu_named_workgroup_barrier_t", AMDGPUNamedWorkgroupBarrier, AMDGPUNamedWorkgroupBarrierTy, 128, 32, 0)
 
 #undef AMDGPU_TYPE
 #undef AMDGPU_OPAQUE_PTR_TYPE
 #undef AMDGPU_NAMED_BARRIER_TYPE
+#undef AMDGPU_IMAGE_RSRC_TYPE
diff --git a/clang/include/clang/Basic/Builtins.def b/clang/include/clang/Basic/Builtins.def
index 48437c9397570..a91315680f93f 100644
--- a/clang/include/clang/Basic/Builtins.def
+++ b/clang/include/clang/Basic/Builtins.def
@@ -34,6 +34,7 @@
 //  Q -> target builtin type, followed by a character to distinguish the builtin type
 //    Qa -> AArch64 svcount_t builtin type.
 //    Qb -> AMDGPU __amdgpu_buffer_rsrc_t builtin type.
+//    Qc -> AMDGPU __amdgpu_image_desc_t builtin type.
 //  E -> ext_vector, followed by the number of elements and the base type.
 //  X -> _Complex, followed by the base type.
 //  Y -> ptrdiff_t
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 97c59b2ceec2f..7ba1dfed1c0db 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -12580,6 +12580,10 @@ static QualType DecodeTypeFromStr(const char *&Str, const ASTContext &Context,
       Type = Context.AMDGPUBufferRsrcTy;
       break;
     }
+    case 'c': {
+      Type = Context.AMDGPUImageDescRsrcTy;
+      break;
+    }
     default:
       llvm_unreachable("Unexpected target builtin type");
     }
diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp
index 578d09f7971d6..3f5277ab66f55 100644
--- a/clang/lib/CodeGen/CGDebugInfo.cpp
+++ b/clang/lib/CodeGen/CGDebugInfo.cpp
@@ -1020,6 +1020,14 @@ llvm::DIType *CGDebugInfo::CreateType(const BuiltinType *BT) {
           DBuilder.createBasicType(Name, Width, llvm::dwarf::DW_ATE_unsigned); \
     return SingletonId;                                                        \
   }
+#define AMDGPU_IMAGE_RSRC_TYPE(Name, Id, SingletonId)                          \
+  case BuiltinType::Id: {                                                      \
+    if (!SingletonId)                                                          \
+      SingletonId =                                                            \
+          DBuilder.createForwardDecl(llvm::dwarf::DW_TAG_structure_type, Name, \
+                                     TheCU, TheCU->getFile(), 0);              \
+    return SingletonId;                                                        \
+  }
 #include "clang/Basic/AMDGPUTypes.def"
   case BuiltinType::UChar:
   case BuiltinType::Char_U:
diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp
index 3ffe999d01178..e3e44556ce514 100644
--- a/clang/lib/CodeGen/CodeGenTypes.cpp
+++ b/clang/lib/CodeGen/CodeGenTypes.cpp
@@ -581,6 +581,10 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) {
   case BuiltinType::Id:                                                        \
     return llvm::TargetExtType::get(getLLVMContext(), "amdgcn.named.barrier",  \
                                     {}, {Scope});
+#define AMDGPU_IMAGE_RSRC_TYPE(Name, Id, SingletonId)                          \
+  case BuiltinType::Id:                                                        \
+    return llvm::VectorType::get(llvm::Type::getInt32Ty(getLLVMContext()), 8,  \
+                                 false);
 #include "clang/Basic/AMDGPUTypes.def"
 #define HLSL_INTANGIBLE_TYPE(Name, Id, SingletonId) case BuiltinType::Id:
 #include "clang/Basic/HLSLIntangibleTypes.def"
diff --git a/clang/test/CodeGen/amdgpu-image-rsrc-type-debug-info.c b/clang/test/CodeGen/amdgpu-image-rsrc-type-debug-info.c
new file mode 100644
index 0000000000000..0e42420e26322
--- /dev/null
+++ b/clang/test/CodeGen/amdgpu-image-rsrc-type-debug-info.c
@@ -0,0 +1,18 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn -emit-llvm -o - %s -debug-info-kind=limited | FileCheck %s
+
+// CHECK-LABEL: define dso_local void @test_locals(
+// CHECK-SAME: ) #[[ATTR0:[0-9]+]] !dbg [[DBG6:![0-9]+]] {
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK-NEXT:    [[IMG:%.*]] = alloca <8 x i32>, align 32, addrspace(5)
+// CHECK-NEXT:    [[IMG_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[IMG]] to ptr
+// CHECK-NEXT:      #dbg_declare(ptr addrspace(5) [[IMG]], [[META11:![0-9]+]], !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef), [[META14:![0-9]+]])
+// CHECK-NEXT:    [[TMP0:%.*]] = load <8 x i32>, ptr [[IMG_ASCAST]], align 32, !dbg [[DBG15:![0-9]+]]
+// CHECK-NEXT:    ret void, !dbg [[DBG16:![0-9]+]]
+//
+void test_locals(void) {
+  __amdgpu_image_rsrc_t img;
+  (void)img;
+}
+
diff --git a/clang/test/CodeGenCXX/amdgpu-image-rsrc-typeinfo.cpp b/clang/test/CodeGenCXX/amdgpu-image-rsrc-typeinfo.cpp
new file mode 100644
index 0000000000000..d96cf5f35c5b2
--- /dev/null
+++ b/clang/test/CodeGenCXX/amdgpu-image-rsrc-typeinfo.cpp
@@ -0,0 +1,7 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn %s -emit-llvm -o - | FileCheck %s
+namespace std { class type_info; }
+auto &a = typeid(__amdgpu_image_rsrc_t);
+//// NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
+// CHECK: {{.*}}
diff --git a/clang/test/SemaCXX/amdgpu-image-rsrc.cpp b/clang/test/SemaCXX/amdgpu-image-rsrc.cpp
new file mode 100644
index 0000000000000..1a19a94039b5e
--- /dev/null
+++ b/clang/test/SemaCXX/amdgpu-image-rsrc.cpp
@@ -0,0 +1,21 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -fsyntax-only -verify -std=gnu++11 -triple amdgcn -Wno-unused-value %s
+
+void foo() {
+  int n = 1;
+  __amdgpu_image_rsrc_t v = 0;            // expected-error {{cannot initialize a variable of type '__amdgpu_image_rsrc_t' with an rvalue of type 'int'}}
+  static_cast<__amdgpu_image_rsrc_t>(n); // expected-error {{static_cast from 'int' to '__amdgpu_image_rsrc_t' is not allowed}}
+  reinterpret_cast<__amdgpu_image_rsrc_t>(n); // expected-error {{reinterpret_cast from 'int' to '__amdgpu_image_rsrc_t' is not allowed}}
+  (void)(v + v); // expected-error {{invalid operands}}
+  int x(v);      // expected-error {{cannot initialize a variable of type 'int' with an lvalue of type '__amdgpu_image_rsrc_t'}}
+  __amdgpu_image_rsrc_t k;
+}
+
+static_assert(sizeof(__amdgpu_image_rsrc_t) == 32, "size");
+static_assert(alignof(__amdgpu_image_rsrc_t) == 32, "align");
+
+template<class T> void bar(T);
+void use(__amdgpu_image_rsrc_t r) { bar(r); }
+struct S { __amdgpu_image_rsrc_t r; int a; };
+static_assert(sizeof(S) == 64, "struct layout");
diff --git a/clang/test/SemaOpenCL/amdgpu-image-rsrc.cl b/clang/test/SemaOpenCL/amdgpu-image-rsrc.cl
new file mode 100644
index 0000000000000..341ab667ebd06
--- /dev/null
+++ b/clang/test/SemaOpenCL/amdgpu-image-rsrc.cl
@@ -0,0 +1,14 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -verify -cl-std=CL1.2 -triple amdgcn-amd-amdhsa %s
+// RUN: %clang_cc1 -verify -cl-std=CL2.0 -triple amdgcn-amd-amdhsa %s
+
+void f() {
+    int n = 3;
+    __amdgpu_image_rsrc_t v = 0; // expected-error {{initializing '__private __amdgpu_image_rsrc_t' with an expression of incompatible type 'int'}}
+    int k = v;                   // expected-error {{initializing '__private int' with an expression of incompatible type '__private __amdgpu_image_rsrc_t'}}
+    (void)(v + v);               // expected-error {{invalid operands}}
+    __amdgpu_image_rsrc_t r;
+    int *p = (int*)r;            // expected-error {{operand of type '__amdgpu_image_rsrc_t' where arithmetic or pointer type is required}}
+    (void)p;
+}
diff --git a/clang/test/SemaOpenMP/amdgpu-image-rsrc.cpp b/clang/test/SemaOpenMP/amdgpu-image-rsrc.cpp
new file mode 100644
index 0000000000000..91d566be9b8a3
--- /dev/null
+++ b/clang/test/SemaOpenMP/amdgpu-image-rsrc.cpp
@@ -0,0 +1,12 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -Wno-unused-value %s
+
+void foo() {
+#pragma omp target
+  {
+    int n = 5;
+    __amdgpu_image_rsrc_t v = 0; // expected-error {{cannot initialize a variable of type '__amdgpu_image_rsrc_t' with an rvalue of type 'int'}}
+    (void)(v + v);               // expected-error {{invalid operands to binary expression ('__amdgpu_image_rsrc_t' and '__amdgpu_image_rsrc_t'}}
+  }
+}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index afce1fe6af854..d41ce5b64b7cd 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -967,6 +967,7 @@ class AMDGPUDimProps<bits<3> enc, string name, string asmsuffix,
   bits<8> NumGradients = !size(GradientArgs);
 }
 
+def AMDGPUImageDescRsrcTy : LLVMType<v8i32>;
 def AMDGPUDim1D : AMDGPUDimProps<0x0, "1d", "1D", ["s"], []>;
 def AMDGPUDim2D : AMDGPUDimProps<0x1, "2d", "2D", ["s", "t"], []>;
 def AMDGPUDim3D : AMDGPUDimProps<0x2, "3d", "3D", ["s", "t", "r"], []>;

clang/include/clang/Basic/Builtins.def

ranapratap55 · 2025-09-24T09:07:19Z

@arsenm ping.

ranapratap55 · 2025-09-24T10:58:33Z

@yxsamliu @jayfoad ping.

clang/include/clang/Basic/AMDGPUTypes.def

clang/include/clang/Basic/Builtins.def

clang/lib/CodeGen/CGDebugInfo.cpp

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

…atch)

clang/lib/CodeGen/CGDebugInfo.cpp

clang/include/clang/Basic/Builtins.def

clang/include/clang/Basic/AMDGPUTypes.def

clang/test/SemaCXX/amdgpu-image-rsrc.cpp

shiltian · 2025-09-25T14:35:32Z

If it's just a struct, why add a built-in type? I don't see the need here. We have the buffer resource type because we use it as a regular pointer in the IR. For the Clang builtins, they can just use v8i32.

arsenm · 2025-09-25T14:53:08Z

If it's just a struct, why add a built-in type? I don't see the need here. We have the buffer resource type because we use it as a regular pointer in the IR. For the Clang builtins, they can just use v8i32.

No, this should be treated as an opaque type. I do not want vector types escaping into code, eventually these should act like pointers too

shiltian · 2025-09-25T15:03:47Z

But @ranapratap55 said it is not a pointer.

ranapratap55 · 2025-09-26T05:47:58Z

But @ranapratap55 said it is not a pointer.

%tmp = tail call float @llvm.amdgcn.image.load.2d.f32.i32(i32 noundef 1, i32 %arg1, i32 %arg2, <8 x i32> %arg3, i32 noundef 0, i32 noundef 0)

intrinsic accepts texture type as <8 x i32>, so in the builtin def also I have defined as v8i32.

ranapratap55 · 2025-09-26T06:03:54Z

If I use opaque ptr macro, getting below IR where the intrinsic texture type will become ptr

%2 = load ptr, ptr %vec8i32.addr.ascast, align 32
%3 = call float @llvm.amdgcn.image.load.2d.f32.i32.p0(i32 12, i32 %0, i32 %1, ptr %2, i32 106, i32 103)

and
If I use new macro, IR generates like below with intrinsic getting the correct <8 x i32>

%2 = load <8 x i32>, ptr %vec8i32.addr.ascast, align 32
%3 = call float @llvm.amdgcn.image.load.2d.f32.i32.v8i32(i32 12, i32 %0, i32 %1, <8 x i32> %2, i32 106, i32 103)

I believe using new macro is better, let me know If I missed anything.

…nor changes

ranapratap55 · 2025-09-26T19:23:35Z

Updated the macro to use opaque ptr type and lowering to <8 x i32> will be in this PR #140210

ranapratap55 · 2025-09-29T08:21:23Z

@arsenm @shiltian @yxsamliu ping.

Adding a new builtin type for AMDGPU's image descriptor rsrc data type This requires for llvm#140210

ranapratap55 requested review from arsenm and shiltian September 23, 2025 09:19

llvmbot added clang Clang issues not falling into any other category backend:AMDGPU clang:frontend Language frontend issues, e.g. anything involving "Sema" clang:codegen IR generation bugs: mangling, exceptions, etc. debuginfo llvm:ir labels Sep 23, 2025

ranapratap55 mentioned this pull request Sep 23, 2025

[AMDGPU] Support for type inferring image load/store builtins for AMDGPU #140210

Open

arsenm requested a review from yxsamliu September 23, 2025 09:21

arsenm reviewed Sep 23, 2025

View reviewed changes

clang/include/clang/Basic/Builtins.def Outdated Show resolved Hide resolved

ranapratap55 added 2 commits September 24, 2025 13:02

[AMDGPU] Add a new builtin type for image descriptor rsrc

2da2208

[AMDGPU] Change image desc rsrc character from 'c' to 't'

2ceded1

ranapratap55 force-pushed the ranapratap55/image-desc-rsrc-builtin-type branch from 5998a77 to 2ceded1 Compare September 24, 2025 07:35

shiltian reviewed Sep 24, 2025

View reviewed changes

clang/include/clang/Basic/AMDGPUTypes.def Outdated Show resolved Hide resolved

shiltian reviewed Sep 24, 2025

View reviewed changes

clang/include/clang/Basic/Builtins.def Outdated Show resolved Hide resolved

shiltian reviewed Sep 24, 2025

View reviewed changes

clang/lib/CodeGen/CGDebugInfo.cpp Outdated Show resolved Hide resolved

arsenm reviewed Sep 25, 2025

View reviewed changes

llvm/include/llvm/IR/IntrinsicsAMDGPU.td Outdated Show resolved Hide resolved

[AMDGPU] Removed def in intrinsic(not related to this clang builtin p…

ef4abb9

…atch)

arsenm reviewed Sep 25, 2025

View reviewed changes

[AMDGPU] Changes image desc macro to use opaque ptr type macro and mi…

f46abee

…nor changes

arsenm approved these changes Sep 29, 2025

View reviewed changes

yxsamliu approved these changes Sep 29, 2025

View reviewed changes

ranapratap55 merged commit 27fa1d0 into llvm:main Sep 29, 2025
9 checks passed

mahesh-attarde pushed a commit to mahesh-attarde/llvm-project that referenced this pull request Oct 3, 2025

[AMDGPU] Add a new builtin type for image descriptor rsrc (llvm#160258)

471f1d0

Adding a new builtin type for AMDGPU's image descriptor rsrc data type This requires for llvm#140210

[AMDGPU] Add a new builtin type for image descriptor rsrc #160258

[AMDGPU] Add a new builtin type for image descriptor rsrc #160258

Uh oh!

Conversation

ranapratap55 commented Sep 23, 2025

Uh oh!

llvmbot commented Sep 23, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Sep 23, 2025

Uh oh!

Uh oh!

ranapratap55 commented Sep 24, 2025

Uh oh!

ranapratap55 commented Sep 24, 2025

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

shiltian commented Sep 25, 2025

Uh oh!

arsenm commented Sep 25, 2025

Uh oh!

shiltian commented Sep 25, 2025

Uh oh!

ranapratap55 commented Sep 26, 2025

Uh oh!

ranapratap55 commented Sep 26, 2025

Uh oh!

ranapratap55 commented Sep 26, 2025

Uh oh!

ranapratap55 commented Sep 29, 2025

Uh oh!

Uh oh!

Uh oh!

llvmbot commented Sep 23, 2025 •

edited

Loading